#include "kernel_operator.h"

using namespace AscendC;
constexpr uint32_t BUFFER_NUM = 2;

class KernelIsCloseFloat
{
    using T = float;

public:
    __aicore__ inline KernelIsCloseFloat() {}
    __aicore__ inline void Init(GM_ADDR x1, GM_ADDR x2, GM_ADDR y,
                                float rtol, float atol, bool equal_nan,
                                uint32_t formerLength, uint32_t formerNum,
                                uint32_t tailLength,
                                uint32_t totalLength, uint32_t partitionLength,
                                TPipe *pipeIn)
    {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        this->rtol = rtol;
        this->atol = atol;
        this->equal_nan = equal_nan;
        const uint32_t weight = sizeof(T);
        this->formerLength = formerLength;
        this->formerNum = formerNum;
        this->tailLength = tailLength;
        this->totalLength = totalLength;

        x1Gm.SetGlobalBuffer((__gm__ T *)x1, totalLength);
        x2Gm.SetGlobalBuffer((__gm__ T *)x2, totalLength);
        yGm.SetGlobalBuffer((__gm__ uint8_t *)y, totalLength);

        pipe = pipeIn;
        pipe->InitBuffer(inQueueX1, BUFFER_NUM, partitionLength * weight * sizeof(T));
        pipe->InitBuffer(inQueueX2, BUFFER_NUM, partitionLength * weight * sizeof(T));
        pipe->InitBuffer(outQueueY, BUFFER_NUM, partitionLength * weight * sizeof(uint8_t));
        pipe->InitBuffer(tmp1, partitionLength * weight * sizeof(half));
        pipe->InitBuffer(tmpSelMask, partitionLength * weight * sizeof(uint8_t));
    }
    __aicore__ inline void Process()
    {
        for (uint32_t i = 0; i < this->formerNum; ++i)
        {
            CopyIn(i, this->formerLength);
            Compute(i, this->formerLength);
            CopyOut(i, this->formerLength);
        }
        if (this->tailLength)
        {
            uint32_t tailLength_AC_32 = ((this->tailLength * sizeof(T) + 31) / 32 * 32) / sizeof(T);
            CopyIn(this->formerNum, tailLength_AC_32);
            Compute(this->formerNum, this->tailLength);
            CopyOut(this->formerNum, tailLength_AC_32);
        }
    }

private:
    __aicore__ inline void CopyIn(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<T> x1Local = inQueueX1.AllocTensor<T>();
        LocalTensor<T> x2Local = inQueueX2.AllocTensor<T>();

        DataCopy(x1Local, x1Gm[_progress_ * this->formerLength], _len_);
        DataCopy(x2Local, x2Gm[_progress_ * this->formerLength], _len_);

        inQueueX1.EnQue(x1Local);
        inQueueX2.EnQue(x2Local);
    }
    __aicore__ inline void Compute(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<T> x1Local = inQueueX1.DeQue<T>();
        LocalTensor<T> x2Local = inQueueX2.DeQue<T>();
        LocalTensor<uint8_t> yLocal = outQueueY.AllocTensor<uint8_t>();
        auto p1 = tmp1.Get<half>();
        auto selMask = tmpSelMask.Get<uint8_t>();

        // |input - other|
        Sub(x1Local, x1Local, x2Local, _len_);
        Abs(x1Local, x1Local, _len_);

        // atol + rtol * |other|
        Abs(x2Local, x2Local, _len_);
        Muls(x2Local, x2Local, this->rtol, _len_);
        Adds(x2Local, x2Local, this->atol, _len_);

        // 比较
        Compare(selMask, x1Local, x2Local, AscendC::CMPMODE::LE, (_len_ + 63) / 64 * 64);
        Duplicate(p1, static_cast<half>(1), _len_);
        Select(p1, selMask, p1, static_cast<half>(0), SELMODE::VSEL_TENSOR_SCALAR_MODE, _len_);
        Cast(yLocal, p1, RoundMode::CAST_NONE, _len_);

        outQueueY.EnQue(yLocal);
        inQueueX1.FreeTensor(x1Local);
        inQueueX2.FreeTensor(x2Local);
    }
    __aicore__ inline void CopyOut(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<uint8_t> yLocal = outQueueY.DeQue<uint8_t>();
        DataCopyPad(yGm[_progress_ * this->formerLength], yLocal, {1, _len_, 0, 0, 0});
        outQueueY.FreeTensor(yLocal);
    }

private:
    TPipe *pipe;
    TQue<QuePosition::VECIN, 1> inQueueX1, inQueueX2;
    TQue<QuePosition::VECOUT, 1> outQueueY;
    TBuf<QuePosition::VECCALC> tmp1, tmpSelMask;

    GlobalTensor<T> x1Gm;
    GlobalTensor<T> x2Gm;
    GlobalTensor<uint8_t> yGm;

    float rtol;
    float atol;
    bool equal_nan;

    uint32_t formerLength;
    uint32_t formerNum;
    uint32_t tailLength;
    uint32_t totalLength;
};

class KernelIsCloseHalf
{
    using T = half;

public:
    __aicore__ inline KernelIsCloseHalf() {}
    __aicore__ inline void Init(GM_ADDR x1, GM_ADDR x2, GM_ADDR y,
                                float rtol, float atol, bool equal_nan,
                                uint32_t formerLength, uint32_t formerNum,
                                uint32_t tailLength,
                                uint32_t totalLength, uint32_t partitionLength,
                                TPipe *pipeIn)
    {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        this->rtol = rtol;
        this->atol = atol;
        this->equal_nan = equal_nan;

        const uint32_t weight = sizeof(T);
        this->formerLength = formerLength;
        this->formerNum = formerNum;
        this->tailLength = tailLength;
        this->totalLength = totalLength;

        x1Gm.SetGlobalBuffer((__gm__ T *)x1, totalLength);
        x2Gm.SetGlobalBuffer((__gm__ T *)x2, totalLength);
        yGm.SetGlobalBuffer((__gm__ uint8_t *)y, totalLength);

        pipe = pipeIn;
        pipe->InitBuffer(inQueueX1, BUFFER_NUM, partitionLength * weight * sizeof(T));
        pipe->InitBuffer(inQueueX2, BUFFER_NUM, partitionLength * weight * sizeof(T));
        pipe->InitBuffer(outQueueY, BUFFER_NUM, partitionLength * weight * sizeof(uint8_t));

        pipe->InitBuffer(tmp1, partitionLength * weight * sizeof(half));
        pipe->InitBuffer(tmp2, partitionLength * weight * sizeof(float));
        pipe->InitBuffer(tmpSelMask, partitionLength * weight * sizeof(uint8_t));
    }
    __aicore__ inline void Process()
    {
        for (uint32_t i = 0; i < this->formerNum; ++i)
        {
            CopyIn(i, this->formerLength);
            Compute(i, this->formerLength);
            CopyOut(i, this->formerLength);
        }
        if (this->tailLength)
        {
            uint32_t tailLength_AC_32 = ((this->tailLength * sizeof(T) + 31) / 32 * 32) / sizeof(T);
            CopyIn(this->formerNum, tailLength_AC_32);
            Compute(this->formerNum, this->tailLength);
            CopyOut(this->formerNum, tailLength_AC_32);
        }
    }

private:
    __aicore__ inline void CopyIn(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<T> x1Local = inQueueX1.AllocTensor<T>();
        LocalTensor<T> x2Local = inQueueX2.AllocTensor<T>();

        DataCopy(x1Local, x1Gm[_progress_ * this->formerLength], _len_);
        DataCopy(x2Local, x2Gm[_progress_ * this->formerLength], _len_);

        inQueueX1.EnQue(x1Local);
        inQueueX2.EnQue(x2Local);
    }
    __aicore__ inline void Compute(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<T> x1Local = inQueueX1.DeQue<T>();
        LocalTensor<T> x2Local = inQueueX2.DeQue<T>();
        LocalTensor<uint8_t> yLocal = outQueueY.AllocTensor<uint8_t>();
        auto p1 = tmp1.Get<half>();
        auto p2 = tmp2.Get<float>();
        auto selMask = tmpSelMask.Get<uint8_t>();

        // |input - other|
        Sub(x1Local, x1Local, x2Local, _len_);
        Abs(x1Local, x1Local, _len_);

        // atol + rtol * |other|
        Abs(x2Local, x2Local, _len_);
        Cast(p2, x2Local, RoundMode::CAST_NONE, _len_);
        Muls(p2, p2, this->rtol, _len_);
        Adds(p2, p2, this->atol, _len_);
        Cast(x2Local, p2, RoundMode::CAST_NONE, _len_);

        // 比较
        Compare(selMask, x1Local, x2Local, AscendC::CMPMODE::LE, (_len_ + 255) / 256 * 256);
        Duplicate(p1, static_cast<half>(1), _len_);
        Select(p1, selMask, p1, static_cast<half>(0), SELMODE::VSEL_TENSOR_SCALAR_MODE, _len_);
        Cast(yLocal, p1, RoundMode::CAST_NONE, _len_);

        outQueueY.EnQue(yLocal);
        inQueueX1.FreeTensor(x1Local);
        inQueueX2.FreeTensor(x2Local);
    }
    __aicore__ inline void CopyOut(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<uint8_t> yLocal = outQueueY.DeQue<uint8_t>();
        DataCopyPad(yGm[_progress_ * this->formerLength], yLocal, {1, _len_, 0, 0, 0});
        outQueueY.FreeTensor(yLocal);
    }

private:
    TPipe *pipe;
    TQue<QuePosition::VECIN, 1> inQueueX1, inQueueX2;
    TQue<QuePosition::VECOUT, 1> outQueueY;
    TBuf<QuePosition::VECCALC> tmp1, tmp2, tmpSelMask;

    GlobalTensor<T> x1Gm;
    GlobalTensor<T> x2Gm;
    GlobalTensor<uint8_t> yGm;

    float rtol;
    float atol;
    bool equal_nan;
    uint32_t formerLength;
    uint32_t formerNum;
    uint32_t tailLength;
    uint32_t totalLength;
};

class KernelIsCloseInt32
{
    using T = int32_t;

public:
    __aicore__ inline KernelIsCloseInt32() {}
    __aicore__ inline void Init(GM_ADDR x1, GM_ADDR x2, GM_ADDR y,
                                float rtol, float atol, bool equal_nan,
                                uint32_t formerLength, uint32_t formerNum,
                                uint32_t tailLength,
                                uint32_t totalLength, uint32_t partitionLength,
                                TPipe *pipeIn)
    {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        this->rtol = rtol;
        this->atol = atol;
        this->equal_nan = equal_nan;

        const uint32_t weight = sizeof(T);
        this->formerLength = formerLength;
        this->formerNum = formerNum;
        this->tailLength = tailLength;
        this->totalLength = totalLength;

        x1Gm.SetGlobalBuffer((__gm__ T *)x1, totalLength);
        x2Gm.SetGlobalBuffer((__gm__ T *)x2, totalLength);
        yGm.SetGlobalBuffer((__gm__ uint8_t *)y, totalLength);

        pipe = pipeIn;
        pipe->InitBuffer(inQueueX1, BUFFER_NUM, partitionLength * weight * sizeof(T));
        pipe->InitBuffer(inQueueX2, BUFFER_NUM, partitionLength * weight * sizeof(T));
        pipe->InitBuffer(outQueueY, BUFFER_NUM, partitionLength * weight * sizeof(uint8_t));

        pipe->InitBuffer(tmp1, partitionLength * weight * sizeof(half));
        pipe->InitBuffer(tmp2, partitionLength * weight * sizeof(float));
        pipe->InitBuffer(tmp3, partitionLength * weight * sizeof(float));
        pipe->InitBuffer(tmpSelMask, partitionLength * weight * sizeof(uint8_t));
    }
    __aicore__ inline void Process()
    {
        for (uint32_t i = 0; i < this->formerNum; ++i)
        {
            CopyIn(i, this->formerLength);
            Compute(i, this->formerLength);
            CopyOut(i, this->formerLength);
        }
        if (this->tailLength)
        {
            uint32_t tailLength_AC_32 = ((this->tailLength * sizeof(T) + 31) / 32 * 32) / sizeof(T);
            CopyIn(this->formerNum, tailLength_AC_32);
            Compute(this->formerNum, this->tailLength);
            CopyOut(this->formerNum, tailLength_AC_32);
        }
    }

private:
    __aicore__ inline void CopyIn(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<T> x1Local = inQueueX1.AllocTensor<T>();
        LocalTensor<T> x2Local = inQueueX2.AllocTensor<T>();

        DataCopy(x1Local, x1Gm[_progress_ * this->formerLength], _len_);
        DataCopy(x2Local, x2Gm[_progress_ * this->formerLength], _len_);

        inQueueX1.EnQue(x1Local);
        inQueueX2.EnQue(x2Local);
    }
    __aicore__ inline void Compute(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<T> x1Local = inQueueX1.DeQue<T>();
        LocalTensor<T> x2Local = inQueueX2.DeQue<T>();
        LocalTensor<uint8_t> yLocal = outQueueY.AllocTensor<uint8_t>();
        auto p1 = tmp1.Get<half>();
        auto p2 = tmp2.Get<float>();
        auto p3 = tmp3.Get<float>();
        auto selMask = tmpSelMask.Get<uint8_t>();

        // |input - other|
        Sub(x1Local, x1Local, x2Local, _len_);
        Cast(p2, x1Local, RoundMode::CAST_NONE, _len_);
        Abs(p2, p2, _len_);

        // atol + rtol * |other|
        Cast(p3, x2Local, RoundMode::CAST_NONE, _len_);
        Abs(p3, p3, _len_);
        Muls(p3, p3, this->rtol, _len_);
        Adds(p3, p3, this->atol, _len_);

        // 比较
        Compare(selMask, p2, p3, AscendC::CMPMODE::LE, (_len_ + 255) / 256 * 256);
        Duplicate(p1, static_cast<half>(1), _len_);
        Select(p1, selMask, p1, static_cast<half>(0), SELMODE::VSEL_TENSOR_SCALAR_MODE, _len_);
        Cast(yLocal, p1, RoundMode::CAST_NONE, _len_);

        outQueueY.EnQue(yLocal);
        inQueueX1.FreeTensor(x1Local);
        inQueueX2.FreeTensor(x2Local);
    }
    __aicore__ inline void CopyOut(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<uint8_t> yLocal = outQueueY.DeQue<uint8_t>();
        DataCopyPad(yGm[_progress_ * this->formerLength], yLocal, {1, _len_, 0, 0, 0});
        outQueueY.FreeTensor(yLocal);
    }

private:
    TPipe *pipe;
    TQue<QuePosition::VECIN, 1> inQueueX1, inQueueX2;
    TQue<QuePosition::VECOUT, 1> outQueueY;
    TBuf<QuePosition::VECCALC> tmp1, tmp2, tmp3, tmpSelMask;

    GlobalTensor<T> x1Gm;
    GlobalTensor<T> x2Gm;
    GlobalTensor<uint8_t> yGm;

    float rtol;
    float atol;
    bool equal_nan;
    uint32_t formerLength;
    uint32_t formerNum;
    uint32_t tailLength;
    uint32_t totalLength;
};

class KernelIsCloseUint8
{
    using T = uint8_t;

public:
    __aicore__ inline KernelIsCloseUint8() {}
    __aicore__ inline void Init(GM_ADDR x1, GM_ADDR x2, GM_ADDR y,
                                float rtol, float atol, bool equal_nan,
                                uint32_t formerLength, uint32_t formerNum,
                                uint32_t tailLength,
                                uint32_t totalLength, uint32_t partitionLength,
                                TPipe *pipeIn)
    {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        this->rtol = rtol;
        this->atol = atol;
        this->equal_nan = equal_nan;

        const uint32_t weight = sizeof(T);
        this->formerLength = formerLength;
        this->formerNum = formerNum;
        this->tailLength = tailLength;
        this->totalLength = totalLength;

        x1Gm.SetGlobalBuffer((__gm__ T *)x1, totalLength);
        x2Gm.SetGlobalBuffer((__gm__ T *)x2, totalLength);
        yGm.SetGlobalBuffer((__gm__ uint8_t *)y, totalLength);

        pipe = pipeIn;
        pipe->InitBuffer(inQueueX1, BUFFER_NUM, partitionLength * weight * sizeof(T));
        pipe->InitBuffer(inQueueX2, BUFFER_NUM, partitionLength * weight * sizeof(T));
        pipe->InitBuffer(outQueueY, BUFFER_NUM, partitionLength * weight * sizeof(uint8_t));

        pipe->InitBuffer(tmp1, partitionLength * weight * sizeof(half));
        pipe->InitBuffer(tmp2, partitionLength * weight * sizeof(half));
        pipe->InitBuffer(tmp3, partitionLength * weight * sizeof(float));
        pipe->InitBuffer(tmp4, partitionLength * weight * sizeof(float));
        pipe->InitBuffer(tmpSelMask, partitionLength * weight * sizeof(uint8_t));
    }
    __aicore__ inline void Process()
    {
        for (uint32_t i = 0; i < this->formerNum; ++i)
        {
            CopyIn(i, this->formerLength);
            Compute(i, this->formerLength);
            CopyOut(i, this->formerLength);
        }
        if (this->tailLength)
        {
            uint32_t tailLength_AC_32 = ((this->tailLength * sizeof(T) + 31) / 32 * 32) / sizeof(T);
            CopyIn(this->formerNum, tailLength_AC_32);
            Compute(this->formerNum, this->tailLength);
            CopyOut(this->formerNum, tailLength_AC_32);
        }
    }

private:
    __aicore__ inline void CopyIn(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<T> x1Local = inQueueX1.AllocTensor<T>();
        LocalTensor<T> x2Local = inQueueX2.AllocTensor<T>();

        DataCopy(x1Local, x1Gm[_progress_ * this->formerLength], _len_);
        DataCopy(x2Local, x2Gm[_progress_ * this->formerLength], _len_);

        inQueueX1.EnQue(x1Local);
        inQueueX2.EnQue(x2Local);
    }
    __aicore__ inline void Compute(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<T> x1Local = inQueueX1.DeQue<T>();
        LocalTensor<T> x2Local = inQueueX2.DeQue<T>();
        LocalTensor<uint8_t> yLocal = outQueueY.AllocTensor<uint8_t>();
        auto p1 = tmp1.Get<half>();
        auto p2 = tmp2.Get<half>();
        auto p3 = tmp3.Get<float>();
        auto p4 = tmp4.Get<float>();
        auto selMask = tmpSelMask.Get<uint8_t>();

        Cast(p1, x1Local, RoundMode::CAST_NONE, _len_);
        Cast(p2, x2Local, RoundMode::CAST_NONE, _len_);
        Cast(p3, p1, RoundMode::CAST_NONE, _len_);
        Cast(p4, p2, RoundMode::CAST_NONE, _len_);

        // |input - other|
        Sub(p3, p3, p4, _len_);
        // 以下代码计算 mask 有问题，另外实际测试并未+256
        // CompareScalar(yLocal, p3, static_cast<float>(0), AscendC::CMPMODE::LT, (_len_ + 255) / 256 * 256);
        // auto mask64 = yLocal.ReinterpretCast<uint64_t>();
        // uint64_t m[2] = {mask64.GetValue(0), 0};
        // Adds(p3, p3, static_cast<float>(256), m, 1, {1, 1, 8, 8});
        Abs(p3, p3, _len_);

        // atol + rtol * |other|
        Muls(p4, p4, this->rtol, _len_);
        Adds(p4, p4, this->atol, _len_);

        // 比较
        Compare(selMask, p3, p4, AscendC::CMPMODE::LE, (_len_ + 255) / 256 * 256);
        Duplicate(p1, static_cast<half>(1), _len_);
        Select(p1, selMask, p1, static_cast<half>(0), SELMODE::VSEL_TENSOR_SCALAR_MODE, _len_);
        Cast(yLocal, p1, RoundMode::CAST_NONE, _len_);

        outQueueY.EnQue(yLocal);
        inQueueX1.FreeTensor(x1Local);
        inQueueX2.FreeTensor(x2Local);
    }
    __aicore__ inline void CopyOut(const uint32_t _progress_, const uint32_t _len_)
    {
        LocalTensor<uint8_t> yLocal = outQueueY.DeQue<uint8_t>();
        DataCopy(yGm[_progress_ * this->formerLength], yLocal, _len_);
        outQueueY.FreeTensor(yLocal);
    }

private:
    TPipe *pipe;
    TQue<QuePosition::VECIN, 1> inQueueX1, inQueueX2;
    TQue<QuePosition::VECOUT, 1> outQueueY;
    TBuf<QuePosition::VECCALC> tmp1, tmp2, tmp3, tmp4, tmpSelMask;

    GlobalTensor<T> x1Gm;
    GlobalTensor<T> x2Gm;
    GlobalTensor<uint8_t> yGm;

    float rtol;
    float atol;
    bool equal_nan;
    uint32_t formerLength;
    uint32_t formerNum;
    uint32_t tailLength;
    uint32_t totalLength;
};

template <typename T>
class KernelIsCloseBroadCast
{

public:
    __aicore__ inline KernelIsCloseBroadCast() {}
    __aicore__ inline void Init(GM_ADDR x1, GM_ADDR x2, GM_ADDR y,
                                float rtol, float atol, bool equal_nan,
                                int32_t y_dimensional,
                                int32_t *y_ndarray, int32_t *x1_ndarray, int32_t *x2_ndarray,
                                int32_t *y_sumndarray, int32_t *x1_sumndarray, int32_t *x2_sumndarray)
    {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");

        this->rtol = rtol;
        this->atol = atol;
        this->equal_nan = equal_nan;

        this->y_dimensional = y_dimensional;

        this->y_ndarray = y_ndarray;
        this->x1_ndarray = x1_ndarray;
        this->x2_ndarray = x2_ndarray;
        this->y_sumndarray = y_sumndarray;
        this->x1_sumndarray = x1_sumndarray;
        this->x2_sumndarray = x2_sumndarray;

        x1Gm.SetGlobalBuffer((__gm__ T *)x1, 1);
        x2Gm.SetGlobalBuffer((__gm__ T *)x2, 1);
        yGm.SetGlobalBuffer((__gm__ int8_t *)y, 1);
    }
    __aicore__ inline void Process()
    {
        int dim = this->y_dimensional;
        float x1, x2;

        for (int j = 0; j < this->y_sumndarray[dim]; j++)
        {
            int x1_start = 0, x2_start = 0;
            for (int k = 0; k < dim; k++)
            {
                if (this->x1_ndarray[k] != 1)
                {
                    x1_start += this->x1_sumndarray[k] * (j / this->y_sumndarray[k] % this->y_ndarray[k]);
                }
                if (this->x2_ndarray[k] != 1)
                {
                    x2_start += this->x2_sumndarray[k] * (j / this->y_sumndarray[k] % this->y_ndarray[k]);
                }
            }
            if constexpr (std::is_same_v<T, uint8_t>)
            {
                x1 = static_cast<float>(static_cast<int32_t>(x1Gm.GetValue(x1_start)));
                x2 = static_cast<float>(static_cast<int32_t>(x2Gm.GetValue(x2_start)));
            }
            else
            {
                x1 = static_cast<float>(x1Gm.GetValue(x1_start));
                x2 = static_cast<float>(x2Gm.GetValue(x2_start));
                // float diff = x1 - x2;
                // float absDiff = diff > 0 ? diff : -diff;
                // absDiff <= (this->atol + this->rtol * x2) ? yGm.SetValue(j, (int8_t)1) : yGm.SetValue(j, (int8_t)0);
            }
            float diff = x1 - x2;
            float absDiff = diff > 0 ? diff : -diff;
            absDiff <= (this->atol + this->rtol * x2) ? yGm.SetValue(j, (int8_t)1) : yGm.SetValue(j, (int8_t)0);
        }
    }

private:
    GlobalTensor<T> x1Gm;
    GlobalTensor<T> x2Gm;
    GlobalTensor<int8_t> yGm;

    float rtol;
    float atol;
    bool equal_nan;

    int32_t y_dimensional;
    int32_t *y_ndarray;
    int32_t *x1_ndarray;
    int32_t *x2_ndarray;

    int32_t *y_sumndarray;
    int32_t *x1_sumndarray;
    int32_t *x2_sumndarray;
};

extern "C" __global__ __aicore__ void is_close(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
{
    TPipe pipe;
    GET_TILING_DATA(tiling_data, tiling);
    if (TILING_KEY_IS(1))
    {
        KernelIsCloseFloat op;
        op.Init(x1, x2, y, tiling_data.rtol, tiling_data.atol, tiling_data.equal_nan,
                tiling_data.formerLength, tiling_data.formerNum,
                tiling_data.tailLength,
                tiling_data.totalLength, tiling_data.partitionLength, &pipe);
        op.Process();
    }
    else if (TILING_KEY_IS(2))
    {
        KernelIsCloseHalf op;
        op.Init(x1, x2, y, tiling_data.rtol, tiling_data.atol, tiling_data.equal_nan,
                tiling_data.formerLength, tiling_data.formerNum,
                tiling_data.tailLength,
                tiling_data.totalLength, tiling_data.partitionLength, &pipe);
        op.Process();
    }
    else if (TILING_KEY_IS(3))
    {
        KernelIsCloseInt32 op;
        op.Init(x1, x2, y, tiling_data.rtol, tiling_data.atol, tiling_data.equal_nan,
                tiling_data.formerLength, tiling_data.formerNum,
                tiling_data.tailLength,
                tiling_data.totalLength, tiling_data.partitionLength, &pipe);
        op.Process();
    }
    else if (TILING_KEY_IS(4))
    {
        KernelIsCloseUint8 op;
        op.Init(x1, x2, y, tiling_data.rtol, tiling_data.atol, tiling_data.equal_nan,
                tiling_data.formerLength, tiling_data.formerNum,
                tiling_data.tailLength,
                tiling_data.totalLength, tiling_data.partitionLength, &pipe);
        op.Process();
    }
    else if (TILING_KEY_IS(5))
    {
        KernelIsCloseBroadCast<DTYPE_X1> op;
        op.Init(x1, x2, y, tiling_data.rtol, tiling_data.atol, tiling_data.equal_nan,
                tiling_data.y_dimensional,
                tiling_data.y_ndarray, tiling_data.x1_ndarray, tiling_data.x2_ndarray,
                tiling_data.y_sumndarray, tiling_data.x1_sumndarray, tiling_data.x2_sumndarray);
        op.Process();
    }
}